typedef struct __S_Layer S_Layer;

struct __S_Layer {
    S_Layer*		m_UpLayer;
	S_Layer*		m_DownLayer;

	uint m_Layer_ID;

	uint			m_InNum;
	uint			m_OutNum;

	//-----------------Node------------------
	float**		m_Weigths;		//m_输入数量*m_输出数量
    float**     m_梯度;

	float*		m_Err;
    float*		m_Out;
    float*      m_Bias;
};

void f_gradient(S_Layer* Layers) {
    size_t 当前层ID_X = get_group_id(0);
    size_t 前层ID_Y = get_group_id(1);

    //S_Layer* UpStream = Layers->m_DownLayer;
    Layers->m_梯度[前层ID_Y][当前层ID_X] += Layers->m_Err[前层ID_Y] * Layers->m_Out[前层ID_Y];
    //printf("m_Err %v4f \n", (float4)(
    //Layers->m_梯度[前层ID_Y][当前层ID_X],  Layers->m_Err[前层ID_Y], Layers->m_Out[前层ID_Y], Layers->m_Layer_ID) );
    //printf("Layers ----------- %f \n", Layers->m_梯度[当前层ID_X][前层ID_Y]);
    if(前层ID_Y==0){
        //S_Layer* UpStream = Layers->m_UpLayer;
        //Layers->m_梯度[Layers->m_OutNum-1][UpStream->m_OutNum] = Layers->m_Err[当前层ID_X] * Layers->m_Out[当前层ID_X];
        //printf("m_Err B %v4f \n", (float4)(
        //Layers->m_梯度[Layers->m_OutNum-1][UpStream->m_OutNum],  Layers->m_Err[当前层ID_X], Layers->m_Out[当前层ID_X], Layers->m_Layer_ID) );

        if(当前层ID_X==0){
            
            if(Layers->m_DownLayer){
                size_t size[2] = {Layers->m_OutNum+1, Layers->m_DownLayer->m_OutNum};
                ndrange_t ndrange1 = ndrange_2D(size);  //跳过最后偏置节点
                
                void(^func_SamplingBlock)(void) = ^{ f_gradient(Layers->m_DownLayer); } ;
                int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
                //printf("当前层ID_X %i\n", 前层ID_Y);
            }
            
        }
    }
    

}

void f_gradient_weigths(S_Layer* Layers, float rate) {
    size_t 当前层ID_X = get_group_id(0);
    size_t 前层ID_Y = get_group_id(1);

    Layers->m_梯度[前层ID_Y][当前层ID_X] = Layers->m_Err[前层ID_Y] * Layers->m_Out[前层ID_Y];
    Layers->m_Weigths[前层ID_Y][当前层ID_X] += Layers->m_梯度[前层ID_Y][当前层ID_X] * rate;
    //printf("m_Err %v4f \n", (float4)(
    //Layers->m_Weigths[当前层ID_X][前层ID_Y],  Layers->m_Err[当前层ID_X], Layers->m_Out[当前层ID_X], Layers->m_Layer_ID) );
    //printf("m_Err %v4f \n", (float4)(Layers->m_梯度[当前层ID_X][前层ID_Y], Layers->m_Out[当前层ID_X], Layers->m_Err[当前层ID_X], rate));

    if(前层ID_Y==0){
        //S_Layer* UpStream = Layers->m_UpLayer;
        //Layers->m_梯度[Layers->m_OutNum-1][UpStream->m_OutNum] = Layers->m_Err[当前层ID_X] * Layers->m_Out[当前层ID_X];
        //Layers->m_Weigths[Layers->m_OutNum-1][UpStream->m_OutNum] = Layers->m_梯度[Layers->m_OutNum-1][UpStream->m_OutNum] * rate;
        //printf("m_Err B %v4f \n", (float3)(
        //Layers->m_梯度[Layers->m_OutNum-1][UpStream->m_OutNum],  Layers->m_Err[当前层ID_X], Layers->m_Out[当前层ID_X]) );

        if(当前层ID_X==0){
            
            if(Layers->m_DownLayer){
                size_t size[2] = {Layers->m_OutNum+1, Layers->m_DownLayer->m_OutNum};
                ndrange_t ndrange1 = ndrange_2D(size);  //跳过最后偏置节点
                
                void(^func_SamplingBlock)(void) = ^{ f_gradient(Layers->m_DownLayer); } ;
                int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
                
            }
        }
    }
}


__kernel void f_aicalc_gradient(__global S_Layer* Layers, float rate) {
    size_t groupID = get_global_linear_id();
    //printf("m_DownLayer %i\n", Layers->m_DownLayer);
    if(groupID==0 && Layers->m_DownLayer) {
        //printf("f_aicalc_gradient %i\n", Layers->m_DownLayer->m_OutNum);
        size_t size[2] = {Layers->m_OutNum+1, Layers->m_DownLayer->m_OutNum};
        ndrange_t ndrange1 = ndrange_2D(size);  //跳过最后偏置节点
        
        if(rate != 0.0){
            void(^func_SamplingBlock)(void) = ^{ f_gradient_weigths(Layers->m_DownLayer, rate); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
        }else{
            void(^func_SamplingBlock)(void) = ^{ f_gradient(Layers->m_DownLayer); } ;
            int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
        }
        
    }
    
}



